__attribute__((global, cooperative)) void reduce_sum_of_squares(float *inp, float *out, size_t elems_per_thread)
{
    // 定义共享内存
    __shared__ float sdata[12];
    // 全局索引
    auto thread_idx = threadIdx.z * (blockDim.x * blockDim.y) + (gridDim.x * threadIdx.y + blockIdx.x) * blockDim.x + threadIdx.x;
    // 定义DTE
    tops_dte_ctx_t dte_input[2];
    tops_dte_ctx_t dte_output[2];


    dte_input[0].init();
    dte_input[1].init();
    dte_output[0].init();
    dte_output[1].init();
    
    __private_dte__ tops_dte_ctx_t dte_outfinal;
    outfinal.init();
    // 事件: 输入DMA完成
    tops::event event_input0;
    tops::event event_input1;
    tops::event event_final;

    static const size_t L1_BUFFER_LEN = 1024;
    auto pos_start = thread_idx * elems_per_thread;
    auto pos_end = pos_start + elems_per_thread;

    __valigned__ float l1_buffer_input[L1_BUFFER_LEN];
    __valigned__ float l1_buf_out[L1_BUFFER_LEN];

    tops::mdspan l1_inp(tops::Private, &l1_buffer_input, L1_BUFFER_LEN);
    tops::mdspan l1_out(tops::Private, &l1_buf_out, L1_BUFFER_LEN);
    tops::mdspan l2_inp(tops::Shared, &sdata[thread_idx], 1);
    tops::mdspan l3_out(tops::Global, out, 1);

    float sum_num = 0;

    for (size_t i = pos_start; i < pos_end; i += L1_BUFFER_LEN)
    {
        tops::mdspan l3_inp(tops::Global, inp + i, L1_BUFFER_LEN);

        // L3 to L1
        tops::memcpy(ctx, l1_inp, l3_inp);

        for (size_t j = 0; j < L1_BUFFER_LEN; j += tops::vlength<vfloat>())
        {
            const auto &v_inp = tops::vload<vfloat>(l1_buffer_input + j);
            // 计算平方和
            const auto &v_pow = tops::vmul<vfloat>(v_inp, v_inp);
            for (int k = 0; k < tops::vlength<vfloat>(); k++)
            {
                sum_num += v_pow[k];
            }
        }
    }

    tops::mdspan p(tops::Private, &sum_num, 1);
    event_input0 = tops::memcpy(dte_output[0], l2_inp, p);
    tops::wait(event_input0);

    for (int i = 1; i < 12; i++)
    {
        sdata[0] += sdata[i];
    }

    tops::mdspan p2(tops::Private, &sdata[0], 1);

    event_final = tops::memcpy(dte_final, l3_out, p2);
    tops::wait(event_final);
    // 销毁DTE
    dte_input[0].destroy();
    dte_input[1].destroy();
    dte_output[0].destroy();
    dte_output[1].destroy();
    dte_final.destroy();
}
    void GCU_EUC_NORM(float *__restrict dev_inp, float *__restrict dev_out, const int nr_elems)
    {

        static const size_t blocks = 1;
        static const size_t threads = 12;

        reduce_sum_of_squares<<<blocks, threads>>>(dev_inp, dev_out, nr_elems / (blocks * threads));
    }